/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <boost/function.hpp>
#include <boost/bind.hpp>
#include <boost/thread.hpp>
#include <boost/thread/condition.hpp>
#include <boost/thread/mutex.hpp>
//#include <boost/date_time/posix_time/posix_time.hpp>

#include <cuqu/cuqu.h>
#ifdef CUPRINTF_INCLUDED
#include <utils/cuPrintf.cu>
#endif

//----------------------------------------------------------------------
//----------------------------------------------------------------------

__device__ cuqu::detail::barrier_t barrier = CUQU_INIT_BARRIER;

__global__ void barrier_kernel(int n_calls, float *o, float *i1, float *i2, int n_floats) 
{
    const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
    const unsigned int block_size = blockDim.x*blockDim.y*blockDim.z;
    const unsigned int grid_size = gridDim.y*gridDim.x;
    const unsigned int gid = tid + bid*block_size;

#ifdef CUPRINTF_CU
    cuPrintfRestrict(0, CUPRINTF_UNRESTRICTED);
    //cuPrintfRestrict(CUPRINTF_UNRESTRICTED, 0);
    //cuPrintf("num_events=%d\n", num_events);
#endif

    int n = 0;
    while(n < n_calls) {
        if(gid < n_floats)
            o[gid] = i1[gid]*n + i2[gid]*(n_calls-n);
        barrier_wait(&barrier, grid_size);
        ++n;
    }
}

//----------------------------------------------------------------------

void test_barrier_kernel(int n_threads, int n_blocks, int n_repeats, int n_calls, bool verbose, bool partial_dump, bool dump)
{
#ifdef CUPRINTF_CU
    cudaPrintfInit();
#endif
    dim3 dimBlock(n_threads, 1, 1);
    dim3 dimGrid(n_blocks, 1, 1);
    cudaEvent_t startEvent, stopEvent;

    if(verbose) printf("# barrier_kernel config: n_threads=%d n_blocks=%d n_repeats=%d n_calls=%d\n", n_threads, n_blocks, n_repeats, n_calls);

    cuqu::detail::barrier_t barrier = CUQU_INIT_BARRIER;
    cuquSafeCall(cudaMemcpyToSymbol("barrier", &barrier, sizeof(barrier), 0));

    cuquSafeCall(cudaEventCreate(&startEvent));
    cuquSafeCall(cudaEventCreate(&stopEvent));

    const int n_floats = n_blocks*n_threads;
    float *d_out, *d_in1, *d_in2;
    cuquSafeCall(cudaMalloc(&d_out,    sizeof(float)*n_floats));
    cuquSafeCall(cudaMalloc(&d_in1,    sizeof(float)*n_floats));
    cuquSafeCall(cudaMemset( d_in1, 0, sizeof(float)*n_floats));
    cuquSafeCall(cudaMalloc(&d_in2,    sizeof(float)*n_floats));
    cuquSafeCall(cudaMemset( d_in2, 0, sizeof(float)*n_floats));
    if(verbose) printf("# allocated 3 GPU buffers of %d floats\n", n_floats);

    float rep_us[n_repeats];
    float rep_ms[n_repeats];

    if(verbose) printf("# n_threads\tn_blocks\ttot tm(ms)\tsigma\t\ttm per call(us)\tsigma\n");

    for(int nr=0; nr<n_repeats; ++nr) {
        // start timer
        cuquSafeCall(cudaEventRecord(startEvent, 0));
        
        // spawn GPU kernel
        barrier_kernel<<<dimGrid, dimBlock>>>(n_calls, d_out, d_in1, d_in2, n_floats);
        
        // stop and sync timer
        cuquSafeCall(cudaEventRecord(stopEvent, 0));
        cuquSafeCall(cudaEventSynchronize(stopEvent));
        
        float milliseconds = 0.0f;
        cuquSafeCall(cudaEventElapsedTime(&milliseconds, startEvent, stopEvent));
        float uspe = milliseconds / n_calls * 1000.0f;
        if(partial_dump) printf("%d\t\t%d\t\t%f\t%f\t%f\t%f\n", n_threads, n_blocks, milliseconds, 0.0, uspe, 0.0);

        rep_us[nr] = uspe;
        rep_ms[nr] = milliseconds;
    }
    // calc average
    double ms = 0,  us = 0;
    double ms2 = 0, us2 = 0;
    double ems = 0, eus = 0;
    double sms = 0, sus = 0;
    for(int nr=0; nr < n_repeats; ++nr) {
        ms  += rep_ms[nr];
        ms2 += rep_ms[nr]*rep_ms[nr];
        us  += rep_us[nr];
        us2 += rep_us[nr]*rep_us[nr];
    }
    ms /= n_repeats;
    us /= n_repeats;
    for(int nr=0; nr < n_repeats; ++nr) {
        sms = (rep_ms[nr] - ms)*(rep_ms[nr] - ms);
        sus = (rep_us[nr] - us)*(rep_us[nr] - us);
    }
    sms = std::sqrt(sms/n_repeats);
    sus = std::sqrt(sus/n_repeats);
    //printf("ms=%f ms2=%f ems=%f sms=%e\n", ms, ms2, ems, sms);
    //printf("us=%f us2=%f eus=%f sus=%e\n", us, us2, eus, sus);
    if(dump) printf("%d\t\t%d\t\t%f\t%f\t%f\t%f\n", n_threads, n_blocks, ms, sms, us, sus);
    fflush(stdout);

    cuquSafeCall(cudaEventDestroy(startEvent));
    cuquSafeCall(cudaEventDestroy(stopEvent));
    
    cuquSafeCall(cudaFree(d_out));
    cuquSafeCall(cudaFree(d_in1));
    cuquSafeCall(cudaFree(d_in2));
        
#ifdef CUPRINTF_CU
    cudaPrintfDisplay(stdout, true);
    cudaPrintfEnd();
#endif
}

//----------------------------------------------------------------------
//----------------------------------------------------------------------

int main(int argc, char* argv[])
{
    cudaDeviceProp deviceProp;
    int idev;

    //
    int n_calls = 1000;
    int n_threads = 64;
    int n_blocks = 2*14; // default 14 = # SM on C2050
    bool sweep_test = false;
 
    // arg scan
    int opt;
    while ((opt = getopt(argc, argv, "hsb:t:c:")) != -1) {
        switch (opt) {
        case 's':
            sweep_test = true;
            break;
        case 'b':
            n_blocks = atoi(optarg);
            break;
        case 't':
            n_threads = atoi(optarg);
            break;
        case 'c':
            n_calls = atoi(optarg);
            break;
        case 'h':
        default:
            fprintf(stderr, "Usage: %s [s] [-c n_calls] [-b n_blocks] [-t n_threads]\n",
                    argv[0]);
            exit(EXIT_FAILURE);
        }
    }

    // CUDA initialization
    printf("# CUDA init\n"); fflush(stdout);
    idev = 0;
    cuquSafeCall(cudaSetDevice(idev));
    cuquSafeCall(cudaGetDeviceProperties(&deviceProp, idev));
    if(!deviceProp.canMapHostMemory) {
        fprintf(stderr, "Device %d cannot map host memory!\n", idev);
        exit(EXIT_FAILURE);
    }
    cuquSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost));

    printf("# barrier implementation: %s\n", CUQU_BARRIER_STR);

    int warmup_n_calls = 100;
    int n_repeats = 4;

    printf("# WARMUP with %d repetitions of %d barrier calls\n", n_repeats, warmup_n_calls); fflush(stdout);
    test_barrier_kernel(n_threads, n_blocks, n_repeats, warmup_n_calls, false, false, false);

    printf("# %d repetitions of %d barrier calls\n", n_repeats, n_calls); fflush(stdout);
    if(sweep_test) {
        bool verbose = true;
        for(int nt=32; nt<=512; nt *= 2)
            for(int nb=1; nb<=4*14; ++nb) {
                test_barrier_kernel(nt, nb, n_repeats, n_calls, verbose, false, true);
                verbose = false;
            }
    } else {
        test_barrier_kernel(n_threads, n_blocks, n_repeats, n_calls, true, true, true);
    }

    cudaThreadExit();

    return EXIT_SUCCESS;
}

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
